﻿namespace RayDen.RayEngine.IntersectionEngines.OpenCL {
    internal partial  class Kernels
    {

        #region Multiple Primitives Per Node

        public static string MP_PbrtBVHKernel = @"
typedef struct {
	float x, y, z;
} Point;

typedef struct {
	float x, y, z;
} Vector;

typedef struct {
	Point o;
	Vector d;
	float mint, maxt;
} Ray;

typedef struct {
	float Distance;
	float U, V; // Barycentric coordinates of the hit point
	unsigned int Index;
} RayHit;

typedef struct {
	unsigned int v[3];
} Triangle;

typedef struct {
	Point pMin, pMax;
} BBox;

typedef struct {
	BBox bbox;
    unsigned int udata;
    unsigned char axis;
    unsigned char count;
    unsigned char b0,b1;
} GpuBVHNode;

void TriangleIntersect(
		const float4 rayOrig,
		const float4 rayDir,
		const float minT,
		float *maxT,
		unsigned int *hitIndex,
		float *hitB1,
		float *hitB2,
		const unsigned int currentIndex,
		__global Point *verts,
		__global Triangle *tris) {

	// Load triangle vertices
	__global Point *p0 = &verts[tris[currentIndex].v[0]];
	__global Point *p1 = &verts[tris[currentIndex].v[1]];
	__global Point *p2 = &verts[tris[currentIndex].v[2]];

	float4 v0 = (float4) (p0->x, p0->y, p0->z, 0.f);
	float4 v1 = (float4) (p1->x, p1->y, p1->z, 0.f);
	float4 v2 = (float4) (p2->x, p2->y, p2->z, 0.f);

	// Calculate intersection
	float4 e1 = v1 - v0;
	float4 e2 = v2 - v0;
	float4 s1 = cross(rayDir, e2);

	const float divisor = dot(s1, e1);
	if (divisor == 0.f)
		return;

	const float invDivisor = 1.f / divisor;

	// Compute first barycentric coordinate
	const float4 d = rayOrig - v0;
	const float b1 = dot(d, s1) * invDivisor;
	if (b1 < 0.f)
		return;

	// Compute second barycentric coordinate
	const float4 s2 = cross(d, e1);
	const float b2 = dot(rayDir, s2) * invDivisor;
	if (b2 < 0.f)
		return;

	const float b0 = 1.f - b1 - b2;
	if (b0 < 0.f)
		return;

	// Compute _t_ to intersection point
	const float t = dot(e2, s2) * invDivisor;
	if (t < minT || t > *maxT)
		return;

	*maxT = t;
	*hitB1 = b1;
	*hitB2 = b2;
	*hitIndex = currentIndex;
}

int BBoxIntersectP(
		const float4 rayOrig, const float4 invRayDir,
		const float mint, const float maxt,
		const float4 pMin, const float4 pMax) {
	const float4 l1 = (pMin - rayOrig) * invRayDir;
	const float4 l2 = (pMax - rayOrig) * invRayDir;
	const float4 tNear = fmin(l1, l2);
	const float4 tFar = fmax(l1, l2);

	float t0 = max(max(max(tNear.x, tNear.y), max(tNear.x, tNear.z)), mint);
    float t1 = min(min(min(tFar.x, tFar.y), min(tFar.x, tFar.z)), maxt);

	return (t1 > t0);
}

int IsDirNegative(const float4 dir, const unsigned char axis)
{
    if (axis == 0)
        return dir.x <0;
    if (axis == 1)
        return dir.y <0;
    if (axis == 2)
        return dir.z <0;
    return 0;
}

__kernel void Intersect(
		__global Ray *rays,
		__global RayHit *rayHits,
		__global Point *verts,
		__global Triangle *tris,
		const unsigned int triangleCount,
		const unsigned int nodeCount,
		__global GpuBVHNode *bvhTree,
		const unsigned int rayCount) {
	// Select the ray to check
	const int gid = get_global_id(0);
	if (gid >= rayCount)
		return;

	float4 rayOrig,rayDir;
	float minT, maxT;
	{
		__global float4 *basePtr =(__global float4 *)&rays[gid];
		float4 data0 = (*basePtr++);
		float4 data1 = (*basePtr);

		rayOrig = (float4)(data0.x, data0.y, data0.z, 0.f);
		rayDir = (float4)(data0.w, data1.x, data1.y, 0.f);

		minT = data1.z;
		maxT = data1.w;
	}

	float4 invRayDir = (float4) 1.f / rayDir;
    int todo[32];
    todo[0] = 0;

	unsigned int hitIndex = 0xffffffffu;
    unsigned int todoOffset = 0;
    float dist = 1e17f;
	unsigned int currentNode = 0; // Root Node
	float b1, b2;
    float u = 0, v = 0;
	unsigned int indexOfHit = 0xffffffffu; // Non-existent

	float4 pMin, pMax, data0, data1;
	__global float4 *basePtr;
	while (true) {
        pMin = (float4)(bvhTree[currentNode].bbox.pMin.x,bvhTree[currentNode].bbox.pMin.y,bvhTree[currentNode].bbox.pMin.z, 0.f);
        pMax = (float4)(bvhTree[currentNode].bbox.pMax.x,bvhTree[currentNode].bbox.pMax.y,bvhTree[currentNode].bbox.pMax.z, 0.f);

        if (currentNode == nodeCount-2) {
            break;
        }
		if (BBoxIntersectP(rayOrig, invRayDir, minT, maxT, pMin, pMax)) 
        {
			const unsigned int triangleIndex = bvhTree[currentNode].udata;
			const unsigned char count = bvhTree[currentNode].count;


			if (count > 0) 
            {
                for(int ti=0;ti<1;ti++)
                {
                    //unsigned int tIndex =  triangleIndex+ti;

				    TriangleIntersect(rayOrig, rayDir, minT, &maxT, &hitIndex, &b1, &b2, triangleIndex, verts, tris);

                    if (maxT < dist) 
                    {
                          dist = maxT;
                          u = b1;
                          v = b2;
                          indexOfHit = hitIndex;
                    }
                }
                if (hitIndex != 0xffffffffu) {
                    break;
                }   
                if( todoOffset == 0 ) {
                        break;
                }
                currentNode = todo[--todoOffset];
            } 
            else 
            {
                const unsigned char axis = bvhTree[currentNode].axis & 0x03;
                if( IsDirNegative(invRayDir, axis)) 
                {
                    todo[todoOffset++] = currentNode + 1;
                    currentNode = triangleIndex;
                }
                else 
                {
                    todo[todoOffset++] = triangleIndex;
                    currentNode = currentNode + 1;
                }

            }
		} 
        else 
        {
			  if( todoOffset == 0 )
              {
                    break;
              }
              currentNode = todo[--todoOffset];
		}
	}

	// Write result
	rayHits[gid].Distance   = dist;
	rayHits[gid].U = u;
	rayHits[gid].V = v;
	rayHits[gid].Index = indexOfHit;
} 
";

        #endregion


        #region Single Primitive Per Node

        public static string PbrtBVHKernel = @"
typedef struct {
	float x, y, z;
} Point;

typedef struct {
	float x, y, z;
} Vector;

typedef struct {
	Point o;
	Vector d;
	float mint, maxt;
} Ray;

typedef struct {
	float Distance;
	float U, V; // Barycentric coordinates of the hit point
	unsigned int Index;
} RayHit;

typedef struct {
	unsigned int v[3];
} Triangle;

typedef struct {
	Point pMin, pMax;
} BBox;

typedef struct {
	BBox bbox;
    unsigned int udata;
    unsigned char bdata;
    unsigned char b0,b1,b2;
} GpuBVHNode;

void TriangleIntersect(
		const float4 rayOrig,
		const float4 rayDir,
		const float minT,
		float *maxT,
		unsigned int *hitIndex,
		float *hitB1,
		float *hitB2,
		const unsigned int currentIndex,
		__global Point *verts,
		__global Triangle *tris) {

	// Load triangle vertices
	__global Point *p0 = &verts[tris[currentIndex].v[0]];
	__global Point *p1 = &verts[tris[currentIndex].v[1]];
	__global Point *p2 = &verts[tris[currentIndex].v[2]];

	float4 v0 = (float4) (p0->x, p0->y, p0->z, 0.f);
	float4 v1 = (float4) (p1->x, p1->y, p1->z, 0.f);
	float4 v2 = (float4) (p2->x, p2->y, p2->z, 0.f);

	// Calculate intersection
	float4 e1 = v1 - v0;
	float4 e2 = v2 - v0;
	float4 s1 = cross(rayDir, e2);

	const float divisor = dot(s1, e1);
	if (divisor == 0.f)
		return;

	const float invDivisor = 1.f / divisor;

	// Compute first barycentric coordinate
	const float4 d = rayOrig - v0;
	const float b1 = dot(d, s1) * invDivisor;
	if (b1 < 0.f)
		return;

	// Compute second barycentric coordinate
	const float4 s2 = cross(d, e1);
	const float b2 = dot(rayDir, s2) * invDivisor;
	if (b2 < 0.f)
		return;

	const float b0 = 1.f - b1 - b2;
	if (b0 < 0.f)
		return;

	// Compute _t_ to intersection point
	const float t = dot(e2, s2) * invDivisor;
	if (t < minT || t > *maxT)
		return;

	*maxT = t;
	*hitB1 = b1;
	*hitB2 = b2;
	*hitIndex = currentIndex;
}

int BBoxIntersectP(
		const float4 rayOrig, const float4 invRayDir,
		const float mint, const float maxt,
		const float4 pMin, const float4 pMax) {
	const float4 l1 = (pMin - rayOrig) * invRayDir;
	const float4 l2 = (pMax - rayOrig) * invRayDir;
	const float4 tNear = fmin(l1, l2);
	const float4 tFar = fmax(l1, l2);

	float t0 = max(max(max(tNear.x, tNear.y), max(tNear.x, tNear.z)), mint);
    float t1 = min(min(min(tFar.x, tFar.y), min(tFar.x, tFar.z)), maxt);

	return (t1 > t0);
}

int IsDirNegative(const float4 dir, const unsigned char axis)
{
    if (axis == 0)
        return dir.x <0;
    if (axis == 1)
        return dir.y <0;
    if (axis == 2)
        return dir.z <0;
    return 0;
}

__kernel void Intersect(
		__global Ray *rays,
		__global RayHit *rayHits,
		__global Point *verts,
		__global Triangle *tris,
		const unsigned int triangleCount,
		const unsigned int nodeCount,
		__global GpuBVHNode *bvhTree,
		const unsigned int rayCount) {
	// Select the ray to check
	const int gid = get_global_id(0);
	if (gid >= rayCount)
		return;

	float4 rayOrig,rayDir;
	float minT, maxT;
	{
		__global float4 *basePtr =(__global float4 *)&rays[gid];
		float4 data0 = (*basePtr++);
		float4 data1 = (*basePtr);

		rayOrig = (float4)(data0.x, data0.y, data0.z, 0.f);
		rayDir = (float4)(data0.w, data1.x, data1.y, 0.f);

		minT = data1.z;
		maxT = data1.w;
	}

	float4 invRayDir = (float4) 1.f / rayDir;
    int todo[64];
	//const unsigned int lid = get_local_id(0);
    todo[0] = 0;

	unsigned int hitIndex = 0xffffffffu;
	//unsigned int hitIndex2 = 0xffffffffu;

    unsigned int todoOffset  = 0;
	unsigned int currentNode = 0; // Root Node
	float b1 = 0.f, b2 = 0.f;

	float4 pMin, pMax, data0, data1;
	__global float4 *basePtr;
	while (true) {
        basePtr =(__global float4 *)&bvhTree[currentNode];
		data0 = (*basePtr++);
		data1 = (*basePtr);

		pMin = (float4)(data0.x, data0.y, data0.z, 0.f);
		pMax = (float4)(data0.w, data1.x, data1.y, 0.f);

        //pMin = (float4)(bvhTree[currentNode].bbox.pMin.x,bvhTree[currentNode].bbox.pMin.y,bvhTree[currentNode].bbox.pMin.z, 0.f);
        //pMax = (float4)(bvhTree[currentNode].bbox.pMax.x,bvhTree[currentNode].bbox.pMax.y,bvhTree[currentNode].bbox.pMax.z, 0.f);

		if (BBoxIntersectP(rayOrig, invRayDir, minT, maxT, pMin, pMax)) 
        {
            const unsigned char isleaf = bvhTree[currentNode].bdata >> 7;
			if (isleaf) 
            {
			    const unsigned int triangleIndex = as_uint(data1.z);
				TriangleIntersect(rayOrig, rayDir, minT, &maxT, &hitIndex, &b1, &b2, triangleIndex, verts, tris);
                
                if( todoOffset == 0 ) {
                        break;
                }
                todoOffset = todoOffset-1;
                currentNode = todo[todoOffset];
            } 
            else 
            {
			    const unsigned int triangleIndex = as_uint(data1.z);;
                const unsigned char axis = bvhTree[currentNode].bdata & 0x03;
                if( IsDirNegative( invRayDir, axis )) 
                {
                    todo[todoOffset] = currentNode + 1;
                    todoOffset = todoOffset + 1;
                    currentNode = triangleIndex;
                }
                else 
                {
                    todo[todoOffset] = triangleIndex;
                    todoOffset = todoOffset+1;
                    currentNode = currentNode + 1;
                }

            }

		} 
        else 
        {
			  if( todoOffset == 0 )
              {
                    break;
              }
              todoOffset = todoOffset-1;
              currentNode = todo[todoOffset];
		}
	}

	// Write result
	rayHits[gid].Distance   = maxT;
	rayHits[gid].U = b1;
	rayHits[gid].V = b2;
	rayHits[gid].Index = hitIndex;
} 
";

        #endregion




        /* TEMP SAVE - WORKING KERNEL
         #region Single Primitive Per Node

        public static string PbrtBVHKernel = @"
typedef struct {
	float x, y, z;
} Point;

typedef struct {
	float x, y, z;
} Vector;

typedef struct {
	Point o;
	Vector d;
	float mint, maxt;
} Ray;

typedef struct {
	float Distance;
	float U, V; // Barycentric coordinates of the hit point
	unsigned int Index;
} RayHit;

typedef struct {
	unsigned int v[3];
} Triangle;

typedef struct {
	Point pMin, pMax;
} BBox;

typedef struct {
	BBox bbox;
    unsigned int udata;
    unsigned char bdata;
    unsigned char b0,b1,b2;
} GpuBVHNode;

void TriangleIntersect(
		const float4 rayOrig,
		const float4 rayDir,
		const float minT,
		float *maxT,
		unsigned int *hitIndex,
		float *hitB1,
		float *hitB2,
		const unsigned int currentIndex,
		__global Point *verts,
		__global Triangle *tris) {

	// Load triangle vertices
	__global Point *p0 = &verts[tris[currentIndex].v[0]];
	__global Point *p1 = &verts[tris[currentIndex].v[1]];
	__global Point *p2 = &verts[tris[currentIndex].v[2]];

	float4 v0 = (float4) (p0->x, p0->y, p0->z, 0.f);
	float4 v1 = (float4) (p1->x, p1->y, p1->z, 0.f);
	float4 v2 = (float4) (p2->x, p2->y, p2->z, 0.f);

	// Calculate intersection
	float4 e1 = v1 - v0;
	float4 e2 = v2 - v0;
	float4 s1 = cross(rayDir, e2);

	const float divisor = dot(s1, e1);
	if (divisor == 0.f)
		return;

	const float invDivisor = 1.f / divisor;

	// Compute first barycentric coordinate
	const float4 d = rayOrig - v0;
	const float b1 = dot(d, s1) * invDivisor;
	if (b1 < 0.f)
		return;

	// Compute second barycentric coordinate
	const float4 s2 = cross(d, e1);
	const float b2 = dot(rayDir, s2) * invDivisor;
	if (b2 < 0.f)
		return;

	const float b0 = 1.f - b1 - b2;
	if (b0 < 0.f)
		return;

	// Compute _t_ to intersection point
	const float t = dot(e2, s2) * invDivisor;
	if (t < minT || t > *maxT)
		return;

	*maxT = t;
	*hitB1 = b1;
	*hitB2 = b2;
	*hitIndex = currentIndex;
}

int BBoxIntersectP(
		const float4 rayOrig, const float4 invRayDir,
		const float mint, const float maxt,
		const float4 pMin, const float4 pMax) {
	const float4 l1 = (pMin - rayOrig) * invRayDir;
	const float4 l2 = (pMax - rayOrig) * invRayDir;
	const float4 tNear = fmin(l1, l2);
	const float4 tFar = fmax(l1, l2);

	float t0 = max(max(max(tNear.x, tNear.y), max(tNear.x, tNear.z)), mint);
    float t1 = min(min(min(tFar.x, tFar.y), min(tFar.x, tFar.z)), maxt);

	return (t1 > t0);
}

int IsDirNegative(const float4 dir, const unsigned char axis)
{
    if (axis == 0)
        return dir.x <0;
    if (axis == 1)
        return dir.y <0;
    if (axis == 2)
        return dir.z <0;
    return 0;
}

__kernel void Intersect(
		__global Ray *rays,
		__global RayHit *rayHits,
		__global Point *verts,
		__global Triangle *tris,
		const unsigned int triangleCount,
		const unsigned int nodeCount,
		__global GpuBVHNode *bvhTree,
		const unsigned int rayCount) {
	// Select the ray to check
	const int gid = get_global_id(0);
	if (gid >= rayCount)
		return;

	float4 rayOrig,rayDir;
	float minT, maxT;
	{
		__global float4 *basePtr =(__global float4 *)&rays[gid];
		float4 data0 = (*basePtr++);
		float4 data1 = (*basePtr);

		rayOrig = (float4)(data0.x, data0.y, data0.z, 0.f);
		rayDir = (float4)(data0.w, data1.x, data1.y, 0.f);

		minT = data1.z;
		maxT = data1.w;
	}

	float4 invRayDir = (float4) 1.f / rayDir;
    __local int todo[64];
    todo[0] = 0;

	unsigned int hitIndex = 0xffffffffu;
	//unsigned int hitIndex2 = 0xffffffffu;

    unsigned int todoOffset = 0;
    //float dist = 1e10f;
	unsigned int currentNode = 0; // Root Node
	float b1 = 0.f, b2 = 0.f;

	float4 pMin, pMax, data0, data1;
	__global float4 *basePtr;
	while (true) {
        pMin = (float4)(bvhTree[currentNode].bbox.pMin.x,bvhTree[currentNode].bbox.pMin.y,bvhTree[currentNode].bbox.pMin.z, 0.f);
        pMax = (float4)(bvhTree[currentNode].bbox.pMax.x,bvhTree[currentNode].bbox.pMax.y,bvhTree[currentNode].bbox.pMax.z, 0.f);

		if (BBoxIntersectP(rayOrig, invRayDir, minT, maxT, pMin, pMax)) 
        {
            const unsigned char isleaf = bvhTree[currentNode].bdata >> 7;
			if (isleaf) 
            {
			    const unsigned int triangleIndex = bvhTree[currentNode].udata;
				TriangleIntersect(rayOrig, rayDir, minT, &maxT, &hitIndex, &b1, &b2, triangleIndex, verts, tris);
                
                if( todoOffset == 0 ) {
                        break;
                }
                todoOffset = todoOffset-1;
                currentNode = todo[todoOffset];
            } 
            else 
            {
			    const unsigned int triangleIndex = bvhTree[currentNode].udata;
                const unsigned char axis = bvhTree[currentNode].bdata & 0x03;
                if( IsDirNegative( invRayDir, axis )) 
                {
                    todo[todoOffset] = currentNode + 1;
                    todoOffset = todoOffset + 1;
                    currentNode = triangleIndex;
                }
                else 
                {
                    todo[todoOffset] = triangleIndex;
                    todoOffset = todoOffset+1;
                    currentNode = currentNode + 1;
                }

            }

		} 
        else 
        {
			  if( todoOffset == 0 )
              {
                    break;
              }
              todoOffset = todoOffset-1;
              currentNode = todo[todoOffset];
		}
	}

	// Write result
	rayHits[gid].Distance   = maxT;
	rayHits[gid].U = b1;
	rayHits[gid].V = b2;
	rayHits[gid].Index = hitIndex;
} 
";

        #endregion
         */
    }
}
